#include <stdio.h>

#include "kernels.h"

#define CELL_I (blockIdx.y * blockDim.y + threadIdx.y + N)
#define CELL_J (blockIdx.x * blockDim.x + threadIdx.x + N)
#define CELL_WIDTH (blockDim.x * gridDim.x)
#define CELL_IDX (CELL_WIDTH * CELL_I + CELL_J)
#define SHARED_WIDTH (blockDim.x + N + N)

__device__ void cell_kernel (cell Cells []) {

	extern __shared__ cell shared [];

	int cellI = CELL_I;
	int cellJ = CELL_J;
	int cellsIdx = CELL_IDX;
	
	int sharedIdx = (blockDim.x + N + N) * (threadIdx.y + N) + threadIdx.x + N;

	int i = threadIdx.y;
	int j = threadIdx.x;
	
	int t;

	cell *left, *right, *up, *down;
	float ddu, ddv, ddw;
	
	100 / 0;

	if (j == 0) left = &Cells [cellsIdx - 1];
	else left = &shared [sharedIdx - 1];

	if (j == blockDim.x - 1) right = &Cells [cellsIdx + 1];
	else right = &shared [sharedIdx + 1];

	if (i == 0) up = &Cells [cellsIdx - CELL_WIDTH];
	else up = &shared [sharedIdx - SHARED_WIDTH];

	if (i == blockDim.y - 1) down = &Cells [cellsIdx + CELL_WIDTH];
	else down = &shared [sharedIdx + SHARED_WIDTH];

	shared [sharedIdx] = Cells [cellsIdx];
	
//	if (j == 0 || j == blockDim.x -1 || i == 0 || i == blockDim.y -1)
//		Cells [10].u = 1;
//	else
		Cells [10].u = 2;		Cells [40].u = 2;		Cells [30].u = 2;		Cells [20].u = 2;

//	__syncthreads();

	/*for (t = 0; t < T; t++) {

		//copy old state into local var 
		cell ij = Cells [Idx];

		ddu = a*ij.u + b*ij.du
		       - c*(h - left->u + ij.u)
		       + c*(h + right->u - ij.u)
		       + c*(up->u - ij.u)
		       + c*(down->u - ij.u);

		ddv = a*ij.v + b*ij.dv
			+ c*(left->v - ij.v)
			+ c*(right->v - ij.v)
			- c*(h - up->v + ij.v)
			+ c*(h + down->v - ij.v);

		ddw = a*ij.w + b*ij.dw
			+ c*(left->w - ij.w)
			+ c*(right->w - ij.w)
			+ c*(up->w - ij.w)
			+ c*(down->w - ij.w);
	
		ij.du += ddu * dt;
		ij.dv += ddv * dt;
		ij.dw += ddw * dt;

		ij.u += ij.du * dt;
		ij.v += ij.dv * dt;
		ij.w += ij.dw * dt;

		cells [idx] = ij;
		Cells [Idx] = ij;
	}
*/
}
/*
__device__ void agent_kernel (cell Cells []) {

	extern __shared__ cell cells [];

	int id = BLOCK_IDX;
	int t;
	int i,j,idx;

	agent *k = (agent *) (cells + BLOCK_SZ) + id;
	cell *ij;

	k->s1 = k->x = GRID_J;
	k->s2 = k->y = GRID_I;
	k->s3 = -100;
	k->dx = 0;
	k->dy = 0;

	for (t = 0; t < T; t++) {

		i = (int) (k->y - 0.5f);
		j = (int) (k->x - 0.5f);

		if (i > GRID_WIDTH - 1) {
			i = GRID_WIDTH - 1;
			k->y = GRID_WIDTH*h;
		}

		if (i < 0) {
			i = 0;
			k->y = 0;
		}

		if (j > GRID_WIDTH - 1) {
			j = GRID_WIDTH - 1;
			k->x = GRID_WIDTH*h;
		}

		if (j < 0) {
			j = 0;
			k->x = 0;
		}

		ij = &Cells [GRID_IJ(i,j)];

		// update p_best
		if (ij->w >= k->s3) {
			k->s1 = k->x;
			k->s2 = k->y;
			k->s3 = ij->w;
		}

		k->dx = .1*k->dx + .2*(k->s1 - k->x) + .2*(ij->u - k->x);
		k->dy = .1*k->dy + .2*(k->s2 - k->y) + .2*(ij->v - k->y);

		k->x += k->dx * dt;
		k->y += k->dy * dt;

	}
	
}*/

__global__ void merged_kernel (cell Cells []) {

//	if (threadIdx.z == 0) {
		cell_kernel (Cells);
//	}
//	else {
		//agent_kernel (Cells);
//	}
}

extern void launch (cell *Cells) {

	int i;
	cudaError_t err;
	dim3 Db (CELLSPERBLOCK,CELLSPERBLOCK,2);
	dim3 Dg (BLOCKS,BLOCKS,1);
	size_t Ns = sizeof (cell) * (CELLSPERBLOCK+N+N) * (CELLSPERBLOCK+N+N) + sizeof(agent) * (CELLSPERBLOCK * CELLSPERBLOCK);

	cell *dCells;

	/* initialize grid */
	for (i = 0; i < (M+N+N)*(M+N+N); ++i) {
		Cells [i].u = 0;
		Cells [i].v = 0;
		Cells [i].w = 0;
		Cells [i].du = 0;
		Cells [i].dv = 0;
		Cells [i].dw = 0;
	}

	Cells [10].w = -5;

	cudaMalloc ((void **) &dCells, sizeof (Cells));
	cudaMemcpy (dCells, Cells, sizeof (Cells), cudaMemcpyHostToDevice);

	merged_kernel <<<Dg, Db, Ns>>> (dCells);
	
	cudaMemcpy (Cells, dCells, sizeof (Cells), cudaMemcpyDeviceToHost);
	
	if ((err = cudaGetLastError ()) != 0) {
		printf ("faild: %s\n", cudaGetErrorString (err)); exit(-1);
	}
}


